/**
 * sum_arrays_timer
 */
#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"

void sumArrays(float *a, float *b, float *res, int size) {
    for (int i = 0; i < size; i += 4) {
        res[i] = a[i] + b[i];
        res[i + 1] = a[i + 1] + b[i + 1];
        res[i + 2] = a[i + 2] + b[i + 2];
        res[i + 3] = a[i + 3] + b[i + 3];
    }
}


__global__ void sumArraysGPU(float *a, float *b, float *res) {
//    int i = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    res[i] = a[i] + b[i];
}

int main(int argc, char **argv) {
    int dev = 0;
    cudaSetDevice(dev);

    int nElem = (1 << 24) + 1;
//    int nElem = 1 << 25;
    printf("Vector size:%d\n", nElem);

    int nByte = sizeof(float) * nElem;

    float *a_h = (float *) malloc(nByte);
    float *b_h = (float *) malloc(nByte);
    float *res_h = (float *) malloc(nByte);

    float *res_from_gpu_h = (float *) malloc(nByte);

    memset(res_h, 0, nByte);
    memset(res_from_gpu_h, 0, nByte);

    float *a_d, *b_d, *res_d;
    CHECK(cudaMalloc((float **) &a_d, nByte));
    CHECK(cudaMalloc((float **) &b_d, nByte));
    CHECK(cudaMalloc((float **) &res_d, nByte));

    initData(a_h, nElem);
    initData(b_h, nElem);

    CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));

    dim3 block(256);
    dim3 grid(nElem / block.x);

    // timer
    double iStart, iElapse;
    iStart = cpuSecond();
    sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d);
    cudaDeviceSynchronize(); //等待核函数执行完成
    iElapse = cpuSecond() - iStart;
    printf("kernel function elapse is %f sec.\n", iElapse);
    printf("Execution configuration<<<%d, %d>>>.\n", grid.x, block.x);

    CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
    sumArrays(a_h, b_h, res_h, nElem);

    checkResult(res_h, res_from_gpu_h, nElem);

    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(res_d);

    free(a_h);
    free(b_h);
    free(res_h);
    free(res_from_gpu_h);

    return 0;
}